#include "kernel_operator.h"
#include <type_traits>
using namespace AscendC;
constexpr int32_t BUFFER_NUM = 1;
template<bool descending> class KernelSort_1 {
public:
    __aicore__ inline KernelSort_1() {}
    __aicore__ inline void Init(GM_ADDR input, GM_ADDR y1, GM_ADDR y2,
                                uint32_t repeat, uint32_t cycles, TPipe* pipeIn) {
        ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
        this->repeat = repeat;
        this->cycles = cycles;
        this->cycles_32 = (this->cycles + 31)/32*32;
        inputGm.SetGlobalBuffer((__gm__ DTYPE_INPUT*)input);
        y1Gm.SetGlobalBuffer((__gm__ DTYPE_INPUT*)y1);
        y2Gm.SetGlobalBuffer((__gm__ int32_t*)y2);
        pipeIn->InitBuffer(inQueueInput, BUFFER_NUM, this->cycles_32 * sizeof(DTYPE_INPUT));
        pipeIn->InitBuffer(outQueueY1, BUFFER_NUM, this->cycles_32 * sizeof(DTYPE_INPUT));
        pipeIn->InitBuffer(outQueueY2, BUFFER_NUM, this->cycles_32 * sizeof(int32_t));
        pipeIn->InitBuffer(QueueTmp1, 6144* 2 * sizeof(float));
        pipeIn->InitBuffer(QueueTmp2, 6144* 2 * sizeof(float));
        pipeIn->InitBuffer(QueueTmp3, 6000 * sizeof(int32_t));
        tmp_index = QueueTmp3.Get<int32_t>();
        {
            tmp_index.SetValue(0, 0);
            tmp_index.SetValue(1, 1);
            tmp_index.SetValue(2, 2);
            tmp_index.SetValue(3, 3);
            tmp_index.SetValue(4, 4);
            tmp_index.SetValue(5, 5);
            tmp_index.SetValue(6, 6);
            tmp_index.SetValue(7, 7);
            Adds(tmp_index[8], tmp_index, (int32_t)8, 8);
            Adds(tmp_index[16], tmp_index, (int32_t)16, 16);
            Adds(tmp_index[32], tmp_index, (int32_t)32, 32);
            Adds(tmp_index[64], tmp_index, (int32_t)64, 64);
            Adds(tmp_index[128], tmp_index, (int32_t)128, 128);
            Adds(tmp_index[256], tmp_index, (int32_t)256, 256);
            Adds(tmp_index[512], tmp_index, (int32_t)512, 512);
            Adds(tmp_index[1024], tmp_index, (int32_t)1024, 1024);
            Adds(tmp_index[2048], tmp_index, (int32_t)2048, 2048);
            Adds(tmp_index[4096], tmp_index, (int32_t)4096, 1904);
        }
    }
    __aicore__ inline void Process() {
        {
            for (int32_t i = 0; i < this->repeat; i++) {
                CopyIn(i);
                Compute(i);
                CopyOut(i);
            }
        }
    }
private:
    __aicore__ inline void CopyIn(int32_t progress) {
        LocalTensor<DTYPE_INPUT> inputLocal = inQueueInput.AllocTensor<DTYPE_INPUT>();
        DataCopy(inputLocal, inputGm[progress * this->cycles], this->cycles_32);
        inQueueInput.EnQue(inputLocal);
    }
    __aicore__ inline void Compute(int32_t progress) {
        LocalTensor<DTYPE_INPUT> inputLocal = inQueueInput.DeQue<DTYPE_INPUT>();
        LocalTensor<DTYPE_INPUT> y1Local = outQueueY1.AllocTensor<DTYPE_INPUT>();
        LocalTensor<int32_t> y2Local = outQueueY2.AllocTensor<int32_t>();
        if constexpr (std::is_same_v<DTYPE_INPUT, half> || std::is_same_v<DTYPE_INPUT, int8_t> || std::is_same_v<DTYPE_INPUT, uint8_t>)
        {
            auto tmp1 = QueueTmp1.Get<half>();
            auto tmp2 = QueueTmp2.Get<half>();
            Duplicate(tmp1, (half)(-65504), 6144);
            if constexpr (std::is_same_v<DTYPE_INPUT, half>)
            {
                Adds(tmp1, inputLocal, (half)0, this->cycles);
            }
            else
            {
                Cast(tmp1, inputLocal, RoundMode::CAST_NONE, this->cycles);
            }
            if constexpr (descending == false)
            {
                Muls(tmp1, tmp1, (half)(-1), this->cycles);
            }
            Sort32<half>(tmp2, tmp1, tmp_index.ReinterpretCast<uint32_t>(), 192);
            MrgSort4Info params;
            params.elementLengths[0] = 32;
            params.elementLengths[1] = 32;
            params.elementLengths[2] = 32;
            params.elementLengths[3] = 32;
            params.ifExhaustedSuspension = false;
            params.validBit = 15;
            params.repeatTimes = 6144/128;
            MrgSortSrcList<half> srcList;
            srcList.src1 = tmp2[0];
            srcList.src2 = tmp2[32*4*1];
            srcList.src3 = tmp2[32*4*2];
            srcList.src4 = tmp2[32*4*3];
            MrgSort(tmp1, srcList, params);
            params.elementLengths[0] = 128;
            params.elementLengths[1] = 128;
            params.elementLengths[2] = 128;
            params.elementLengths[3] = 128;
            params.repeatTimes = 12;
            srcList.src1 = tmp1[0];
            srcList.src2 = tmp1[128*4*1];
            srcList.src3 = tmp1[128*4*2];
            srcList.src4 = tmp1[128*4*3];
            MrgSort(tmp2, srcList, params);
            params.elementLengths[0] = 512;
            params.elementLengths[1] = 512;
            params.elementLengths[2] = 512;
            params.elementLengths[3] = 512;
            params.repeatTimes = 3;
            srcList.src1 = tmp2[0];
            srcList.src2 = tmp2[512*4*1];
            srcList.src3 = tmp2[512*4*2];
            srcList.src4 = tmp2[512*4*3];
            MrgSort(tmp1, srcList, params);
            params.elementLengths[0] = 2048;
            params.elementLengths[1] = 2048;
            params.elementLengths[2] = 2048;
            params.validBit = 7;
            params.repeatTimes = 1;
            srcList.src1 = tmp1[0];
            srcList.src2 = tmp1[2048*4*1];
            srcList.src3 = tmp1[2048*4*2];
            MrgSort(tmp2, srcList, params);
            uint64_t rsvdCnt = 0; 
            GatherMask(y2Local, tmp2.ReinterpretCast<int32_t>(), 2, false, 0, {1, static_cast<uint16_t>(this->cycles_32/32), 8, 0 }, rsvdCnt);
            GatherMask(tmp1, tmp2, 3, false, 0, {1, static_cast<uint16_t>(this->cycles_32/32), 8, 0 }, rsvdCnt);
            if constexpr (descending == false)
            {
                Muls(tmp1, tmp1, (half)(-1), this->cycles);
            }
            if constexpr (std::is_same_v<DTYPE_INPUT, half>)
            {
                Adds(y1Local, tmp1, (half)0, this->cycles);
            }
            else
            {
                Cast(y1Local, tmp1, RoundMode::CAST_RINT, this->cycles);
            }
        }
        else
        {
            auto tmp1 = QueueTmp1.Get<float>();
            auto tmp2 = QueueTmp2.Get<float>();
            Duplicate(tmp1, (float)(-3.4028235e38), 6144);
            if constexpr (std::is_same_v<DTYPE_INPUT, float>)
            {
                Adds(tmp1, inputLocal, (float)0, this->cycles);
            }
            else if constexpr (std::is_same_v<DTYPE_INPUT, int64_t>)
            {
                Cast(tmp1, inputLocal, RoundMode::CAST_RINT, this->cycles);
            }
            else
            {
                Cast(tmp1, inputLocal, RoundMode::CAST_NONE, this->cycles);
            }
            if constexpr (descending == false)
            {
                Muls(tmp1, tmp1, (float)(-1), this->cycles);
            }
            Sort32<float>(tmp2, tmp1, tmp_index.ReinterpretCast<uint32_t>(), 192);
            MrgSort4Info params;
            params.elementLengths[0] = 32;
            params.elementLengths[1] = 32;
            params.elementLengths[2] = 32;
            params.elementLengths[3] = 32;
            params.ifExhaustedSuspension = false;
            params.validBit = 15;
            params.repeatTimes = 6144/128;
            MrgSortSrcList<float> srcList;
            srcList.src1 = tmp2[0];
            srcList.src2 = tmp2[32/2*4*1];
            srcList.src3 = tmp2[32/2*4*2];
            srcList.src4 = tmp2[32/2*4*3];
            MrgSort(tmp1, srcList, params);
            params.elementLengths[0] = 128;
            params.elementLengths[1] = 128;
            params.elementLengths[2] = 128;
            params.elementLengths[3] = 128;
            params.repeatTimes = 12;
            srcList.src1 = tmp1[0];
            srcList.src2 = tmp1[128/2*4*1];
            srcList.src3 = tmp1[128/2*4*2];
            srcList.src4 = tmp1[128/2*4*3];
            MrgSort(tmp2, srcList, params);
            params.elementLengths[0] = 512;
            params.elementLengths[1] = 512;
            params.elementLengths[2] = 512;
            params.elementLengths[3] = 512;
            params.repeatTimes = 3;
            srcList.src1 = tmp2[0];
            srcList.src2 = tmp2[512/2*4*1];
            srcList.src3 = tmp2[512/2*4*2];
            srcList.src4 = tmp2[512/2*4*3];
            MrgSort(tmp1, srcList, params);
            params.elementLengths[0] = 2048;
            params.elementLengths[1] = 2048;
            params.elementLengths[2] = 2048;
            params.validBit = 7;
            params.repeatTimes = 1;
            srcList.src1 = tmp1[0];
            srcList.src2 = tmp1[2048/2*4*1];
            srcList.src3 = tmp1[2048/2*4*2];
            MrgSort(tmp2, srcList, params);
            uint64_t rsvdCnt = 0; 
            GatherMask(y2Local, tmp2.ReinterpretCast<int32_t>(), 2, false, 0, {1, static_cast<uint16_t>(this->cycles_32/32), 8, 0 }, rsvdCnt);
            GatherMask(tmp1, tmp2, 1, false, 0, {1, static_cast<uint16_t>(this->cycles_32/32), 8, 0 }, rsvdCnt);
            if constexpr (descending == false)
            {
                Muls(tmp1, tmp1, (float)(-1), this->cycles);
            }
            if constexpr (std::is_same_v<DTYPE_INPUT, float>)
            {
                Adds(y1Local, tmp1, (float)0, this->cycles);
            }
            else
            {
                Cast(y1Local, tmp1, RoundMode::CAST_RINT, this->cycles);
            }
        }
        inQueueInput.FreeTensor(inputLocal);
        outQueueY1.EnQue<DTYPE_INPUT>(y1Local);
        outQueueY2.EnQue<int32_t>(y2Local);
    }
    __aicore__ inline void CopyOut(int32_t progress) {
        LocalTensor<DTYPE_INPUT> y1Local = outQueueY1.DeQue<DTYPE_INPUT>();
        LocalTensor<int32_t> y2Local = outQueueY2.DeQue<int32_t>();
        DataCopyExtParams copyParams{1, static_cast<uint32_t>(this->cycles * sizeof(DTYPE_INPUT)), 0, 0, 0};
        DataCopyExtParams copyParams1{1, static_cast<uint32_t>(this->cycles * sizeof(int32_t)), 0, 0, 0};
        DataCopyPad(y1Gm[progress * this->cycles], y1Local, copyParams);
        DataCopyPad(y2Gm[progress * this->cycles], y2Local, copyParams1);
        outQueueY1.FreeTensor(y1Local);
        outQueueY2.FreeTensor(y2Local);
    }
private:
    TQue<QuePosition::VECIN, BUFFER_NUM> inQueueInput;
    TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueY1, outQueueY2;
    TBuf<QuePosition::VECCALC> QueueTmp1, QueueTmp2,QueueTmp3,QueueTmp4, QueueTmp5;
    LocalTensor<int32_t> tmp_index;
    GlobalTensor<DTYPE_INPUT> inputGm;
    GlobalTensor<DTYPE_INPUT> y1Gm;
    GlobalTensor<int32_t> y2Gm;
    uint32_t repeat;
    uint32_t cycles;
    uint32_t cycles_32;
};
template<bool descending> class KernelSort {
public:
    __aicore__ inline KernelSort() {}
    __aicore__ inline void Init(GM_ADDR input, GM_ADDR y1, GM_ADDR y2,
                                uint32_t repeat, uint32_t cycles, uint32_t interval, 
                                uint32_t loopCount, uint32_t tileDataLenght, uint32_t tailDataLenght, TPipe* pipeIn) {
        ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
        this->loopCount = loopCount;
        this->tileDataLenght = tileDataLenght;
        this->tailDataLenght = tailDataLenght;
        this->repeat = repeat;
        this->cycles = cycles;
        this->interval = interval;
        this->cycles_32 = (this->cycles + 31)/32*32;
        inputGm.SetGlobalBuffer((__gm__ DTYPE_INPUT*)input);
        y1Gm.SetGlobalBuffer((__gm__ DTYPE_INPUT*)y1);
        y2Gm.SetGlobalBuffer((__gm__ int32_t*)y2);
        pipeIn->InitBuffer(inQueueInput, BUFFER_NUM, this->cycles_32 * 32);
        pipeIn->InitBuffer(outQueueY1, BUFFER_NUM, this->cycles_32 * 32);
        pipeIn->InitBuffer(outQueueY2, BUFFER_NUM, this->cycles_32 * this->tileDataLenght * sizeof(int32_t));
        pipeIn->InitBuffer(QueueTmp1, 2048* 2 * sizeof(float));
        pipeIn->InitBuffer(QueueTmp2, 2048* 2 * sizeof(float));
        pipeIn->InitBuffer(QueueTmp3, 2048 * sizeof(int32_t));
        pipeIn->InitBuffer(QueueTmp4, 2048 * sizeof(uint32_t));
        pipeIn->InitBuffer(QueueTmp5, 2048 * sizeof(uint32_t));
        tmp_index = QueueTmp3.Get<int32_t>();
        tmp_offset = QueueTmp4.Get<uint32_t>();
        {
            tmp_index.SetValue(0, 0);
            tmp_index.SetValue(1, 1);
            tmp_index.SetValue(2, 2);
            tmp_index.SetValue(3, 3);
            tmp_index.SetValue(4, 4);
            tmp_index.SetValue(5, 5);
            tmp_index.SetValue(6, 6);
            tmp_index.SetValue(7, 7);
            Adds(tmp_index[8], tmp_index, (int32_t)8, 8);
            Adds(tmp_index[16], tmp_index, (int32_t)16, 16);
            Adds(tmp_index[32], tmp_index, (int32_t)32, 32);
            Adds(tmp_index[64], tmp_index, (int32_t)64, 64);
            Adds(tmp_index[128], tmp_index, (int32_t)128, 128);
            Adds(tmp_index[256], tmp_index, (int32_t)256, 256);
            Adds(tmp_index[512], tmp_index, (int32_t)512, 512);
            Adds(tmp_index[1024], tmp_index, (int32_t)1024, 1024);
        }
            Muls(tmp_offset.ReinterpretCast<int32_t>(), tmp_index, (int32_t)32, 2048);
            tmp_offset1 = QueueTmp5.Get<uint32_t>();
            Muls(tmp_offset1.ReinterpretCast<int32_t>(), tmp_index, (int32_t)(32/sizeof(DTYPE_INPUT)*4), 2048);
    }
    __aicore__ inline void Process() {
        {
            for (int32_t i = 0; i < this->repeat; i++) 
            {
                this->processDataLenght = this->tileDataLenght;
                for(int32_t j=0; j<this->loopCount; j++)
                {
                    CopyIn(i, j);
                    LocalTensor<DTYPE_INPUT> inputLocal = inQueueInput.DeQue<DTYPE_INPUT>();
                    LocalTensor<DTYPE_INPUT> y1Local = outQueueY1.AllocTensor<DTYPE_INPUT>();
                    LocalTensor<int32_t> y2Local = outQueueY2.AllocTensor<int32_t>();
                    for(int32_t z=0; z<this->processDataLenght; z++)
                    {
                        Compute(inputLocal, y1Local, y2Local, z);
                    }
                    inQueueInput.FreeTensor(inputLocal);
                    outQueueY1.EnQue<DTYPE_INPUT>(y1Local);
                    outQueueY2.EnQue<int32_t>(y2Local);
                    CopyOut(i, j);
                }
                this->processDataLenght = this->tailDataLenght;
                CopyIn(i, this->loopCount);
                LocalTensor<DTYPE_INPUT> inputLocal = inQueueInput.DeQue<DTYPE_INPUT>();
                LocalTensor<DTYPE_INPUT> y1Local = outQueueY1.AllocTensor<DTYPE_INPUT>();
                LocalTensor<int32_t> y2Local = outQueueY2.AllocTensor<int32_t>();
                for(int32_t z=0; z<this->processDataLenght; z++)
                {
                    Compute(inputLocal, y1Local, y2Local, z);
                }
                inQueueInput.FreeTensor(inputLocal);
                outQueueY1.EnQue<DTYPE_INPUT>(y1Local);
                outQueueY2.EnQue<int32_t>(y2Local);
                CopyOut(i, this->loopCount);
            }
        }
    }
private:
    __aicore__ inline void CopyIn(int32_t progress, int32_t offset) {
        LocalTensor<DTYPE_INPUT> inputLocal = inQueueInput.AllocTensor<DTYPE_INPUT>();
        if constexpr (!std::is_same_v<DTYPE_INPUT, int64_t>) {
            DataCopyExtParams copyParams{static_cast<uint16_t>(this->cycles_32), 32, static_cast<uint32_t>(this->interval*sizeof(DTYPE_INPUT)-32), 0, 0}; 
            DataCopyPadExtParams<DTYPE_INPUT> padParams{false, 0, 0, 0};
            DataCopyPad(inputLocal, inputGm[progress * this->cycles * this->interval + offset*this->tileDataLenght], copyParams, padParams); 
        }
        inQueueInput.EnQue(inputLocal);
    }
    __aicore__ inline void Compute(LocalTensor<DTYPE_INPUT> inputLocal, LocalTensor<DTYPE_INPUT> y1Local, LocalTensor<int32_t> y2Local, int32_t progress) {
        if constexpr (std::is_same_v<DTYPE_INPUT, half> || std::is_same_v<DTYPE_INPUT, int8_t> || std::is_same_v<DTYPE_INPUT, uint8_t>)
        {
            auto tmp1 = QueueTmp1.Get<half>();
            auto tmp2 = QueueTmp2.Get<half>();
            Duplicate(tmp1, (half)(-65504), 2048);
            if constexpr (std::is_same_v<DTYPE_INPUT, half>)
            {
                Gather(tmp1, inputLocal, tmp_offset, (uint32_t)progress*sizeof(DTYPE_INPUT), this->cycles);
            }
            else
            {
                Gather(tmp2.ReinterpretCast<DTYPE_INPUT>(), inputLocal, tmp_offset, (uint32_t)progress, this->cycles);
                Cast(tmp1, tmp2.ReinterpretCast<DTYPE_INPUT>(), RoundMode::CAST_NONE, this->cycles);
            }
            if constexpr (descending == false)
            {
                Muls(tmp1, tmp1, (half)(-1), this->cycles);
            }
            Sort32<half>(tmp2, tmp1, tmp_index.ReinterpretCast<uint32_t>(), 64);
            MrgSort4Info params;
            params.elementLengths[0] = 32;
            params.elementLengths[1] = 32;
            params.elementLengths[2] = 32;
            params.elementLengths[3] = 32;
            params.ifExhaustedSuspension = false;
            params.validBit = 15;
            params.repeatTimes = 16;
            MrgSortSrcList<half> srcList;
            srcList.src1 = tmp2[0];
            srcList.src2 = tmp2[32*4*1];
            srcList.src3 = tmp2[32*4*2];
            srcList.src4 = tmp2[32*4*3];
            MrgSort(tmp1, srcList, params);
            params.elementLengths[0] = 128;
            params.elementLengths[1] = 128;
            params.elementLengths[2] = 128;
            params.elementLengths[3] = 128;
            params.repeatTimes = 4;
            srcList.src1 = tmp1[0];
            srcList.src2 = tmp1[128*4*1];
            srcList.src3 = tmp1[128*4*2];
            srcList.src4 = tmp1[128*4*3];
            MrgSort(tmp2, srcList, params);
            params.elementLengths[0] = 512;
            params.elementLengths[1] = 512;
            params.elementLengths[2] = 512;
            params.elementLengths[3] = 512;
            params.repeatTimes = 1;
            srcList.src1 = tmp2[0];
            srcList.src2 = tmp2[512*4*1];
            srcList.src3 = tmp2[512*4*2];
            srcList.src4 = tmp2[512*4*3];
            MrgSort(tmp1, srcList, params);
            uint64_t rsvdCnt = 0; 
            GatherMask(tmp2, tmp1, 3, false, 0, {1, static_cast<uint16_t>(this->cycles_32/32), 8, 0 }, rsvdCnt);
            if constexpr (descending == false)
            {
                Muls(tmp2, tmp2, (half)(-1), this->cycles);
            }
            if constexpr (std::is_same_v<DTYPE_INPUT, half>)
            {
                Scatter(y1Local, tmp2, tmp_offset, (uint32_t)progress*2, this->cycles);
            }
            else
            {
                Cast(tmp2.ReinterpretCast<DTYPE_INPUT>(), tmp2, RoundMode::CAST_RINT, this->cycles);
                Scatter(y1Local, tmp2.ReinterpretCast<DTYPE_INPUT>(), tmp_offset, (uint32_t)progress*sizeof(DTYPE_INPUT), this->cycles);
            }
            GatherMask(tmp2.ReinterpretCast<int32_t>(), tmp1.ReinterpretCast<int32_t>(), 2, false, 0, {1, static_cast<uint16_t>(this->cycles_32/32), 8, 0 }, rsvdCnt);
            Scatter(y2Local, tmp2.ReinterpretCast<int32_t>(), tmp_offset1, (uint32_t)progress*4, this->cycles);
        }
        else
        {
            auto tmp1 = QueueTmp1.Get<float>();
            auto tmp2 = QueueTmp2.Get<float>();
            Duplicate(tmp1, (float)(-3.4028235e38), 2048);
            if constexpr (std::is_same_v<DTYPE_INPUT, float>)
            {
                Gather(tmp1, inputLocal, tmp_offset, (uint32_t)progress*sizeof(DTYPE_INPUT), this->cycles);
            }
            else
            {
                if constexpr (!std::is_same_v<DTYPE_INPUT, int64_t>) 
                {
                    Gather(tmp2.ReinterpretCast<DTYPE_INPUT>(), inputLocal, tmp_offset, (uint32_t)progress*sizeof(DTYPE_INPUT), this->cycles);
                }
                Cast(tmp1, tmp2.ReinterpretCast<DTYPE_INPUT>(), RoundMode::CAST_NONE, this->cycles);
            }
            if constexpr (descending == false)
            {
                Muls(tmp1, tmp1, (float)(-1), this->cycles);
            }
            Sort32<float>(tmp2, tmp1, tmp_index.ReinterpretCast<uint32_t>(), 64);
            MrgSort4Info params;
            params.elementLengths[0] = 32;
            params.elementLengths[1] = 32;
            params.elementLengths[2] = 32;
            params.elementLengths[3] = 32;
            params.ifExhaustedSuspension = false;
            params.validBit = 15;
            params.repeatTimes = 16;
            MrgSortSrcList<float> srcList;
            srcList.src1 = tmp2[0];
            srcList.src2 = tmp2[32/2*4*1];
            srcList.src3 = tmp2[32/2*4*2];
            srcList.src4 = tmp2[32/2*4*3];
            MrgSort(tmp1, srcList, params);
            params.elementLengths[0] = 128;
            params.elementLengths[1] = 128;
            params.elementLengths[2] = 128;
            params.elementLengths[3] = 128;
            params.repeatTimes = 4;
            srcList.src1 = tmp1[0];
            srcList.src2 = tmp1[128/2*4*1];
            srcList.src3 = tmp1[128/2*4*2];
            srcList.src4 = tmp1[128/2*4*3];
            MrgSort(tmp2, srcList, params);
            params.elementLengths[0] = 512;
            params.elementLengths[1] = 512;
            params.elementLengths[2] = 512;
            params.elementLengths[3] = 512;
            params.repeatTimes = 1;
            srcList.src1 = tmp2[0];
            srcList.src2 = tmp2[512/2*4*1];
            srcList.src3 = tmp2[512/2*4*2];
            srcList.src4 = tmp2[512/2*4*3];
            MrgSort(tmp1, srcList, params);
            uint64_t rsvdCnt = 0; 
            GatherMask(tmp2, tmp1, 1, false, 0, {1, static_cast<uint16_t>(this->cycles_32/32), 8, 0 }, rsvdCnt);
            if constexpr (descending == false)
            {
                Muls(tmp2, tmp2, (float)(-1), this->cycles);
            }
            if constexpr (std::is_same_v<DTYPE_INPUT, float>)
            {
                Scatter(y1Local, tmp2, tmp_offset, (uint32_t)progress*sizeof(DTYPE_INPUT), this->cycles);
            }
            else
            {
                Cast(tmp2.ReinterpretCast<DTYPE_INPUT>(), tmp2, RoundMode::CAST_RINT, this->cycles);
                if constexpr (!std::is_same_v<DTYPE_INPUT, int64_t>) {
                Scatter(y1Local, tmp2.ReinterpretCast<DTYPE_INPUT>(), tmp_offset, (uint32_t)progress*sizeof(DTYPE_INPUT), this->cycles);
                }
            }
            GatherMask(tmp2.ReinterpretCast<int32_t>(), tmp1.ReinterpretCast<int32_t>(), 2, false, 0, {1, static_cast<uint16_t>(this->cycles_32/32), 8, 0 }, rsvdCnt);
            Scatter(y2Local, tmp2.ReinterpretCast<int32_t>(), tmp_offset1, (uint32_t)progress*4, this->cycles);
        }
    }
    __aicore__ inline void CopyOut(int32_t progress, int32_t offset) {
        DataCopyExtParams copyParams1{static_cast<uint16_t>(this->cycles), static_cast<uint32_t>(this->processDataLenght*sizeof(DTYPE_INPUT)), 0, static_cast<uint32_t>(this->interval*sizeof(DTYPE_INPUT)-this->processDataLenght*sizeof(DTYPE_INPUT)), 0}; 
        DataCopyExtParams copyParams2{static_cast<uint16_t>(this->cycles), static_cast<uint32_t>(this->processDataLenght*4), 0, static_cast<uint32_t>(this->interval*4-this->processDataLenght*4), 0}; 
        LocalTensor<DTYPE_INPUT> y1Local = outQueueY1.DeQue<DTYPE_INPUT>();
        LocalTensor<int32_t> y2Local = outQueueY2.DeQue<int32_t>();
        DataCopyPad(y1Gm[progress * this->cycles * this->interval + offset*this->tileDataLenght], y1Local, copyParams1);
        DataCopyPad(y2Gm[progress * this->cycles * this->interval + offset*this->tileDataLenght], y2Local, copyParams2);
        outQueueY1.FreeTensor(y1Local);
        outQueueY2.FreeTensor(y2Local);
    }
private:
    TQue<QuePosition::VECIN, BUFFER_NUM> inQueueInput;
    TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueY1, outQueueY2;
    TBuf<QuePosition::VECCALC> QueueTmp1, QueueTmp2,QueueTmp3,QueueTmp4, QueueTmp5;
    LocalTensor<int32_t> tmp_index;
    LocalTensor<uint32_t> tmp_offset,tmp_offset1;
    GlobalTensor<DTYPE_INPUT> inputGm;
    GlobalTensor<DTYPE_INPUT> y1Gm;
    GlobalTensor<int32_t> y2Gm;
    uint32_t repeat;
    uint32_t cycles;
    uint32_t cycles_32;
    uint32_t interval;
    uint32_t loopCount;
    uint32_t tileDataLenght;
    uint32_t tailDataLenght;
    uint32_t processDataLenght;
};
extern "C" __global__ __aicore__ void sort(GM_ADDR input, GM_ADDR y1, GM_ADDR y2, GM_ADDR workspace, GM_ADDR tiling) {
    GET_TILING_DATA(tiling_data, tiling);
    TPipe pipe;
    if(TILING_KEY_IS(1)) { 
        KernelSort_1<false> op;
        op.Init(input, y1, y2,
                tiling_data.repeat, tiling_data.cycles, &pipe);
        op.Process();
    }
    else if(TILING_KEY_IS(2)) { 
        KernelSort_1<true> op;
        op.Init(input, y1, y2,
                tiling_data.repeat, tiling_data.cycles, &pipe);
        op.Process();
    }
    else if(TILING_KEY_IS(3)) { 
        KernelSort<false> op;
        op.Init(input, y1, y2,
                tiling_data.repeat, tiling_data.cycles, tiling_data.interval, 
                tiling_data.loopCount, tiling_data.tileDataLenght, tiling_data.tailDataLenght, &pipe);
        op.Process();
    }
    else if(TILING_KEY_IS(4)) { 
        KernelSort<true> op;
        op.Init(input, y1, y2,
                tiling_data.repeat, tiling_data.cycles, tiling_data.interval, 
                tiling_data.loopCount, tiling_data.tileDataLenght, tiling_data.tailDataLenght, &pipe);
        op.Process();
    }
}